#include <stdio.h>
#include <sys/time.h>
#include "FreeImage.h"

// SDKBASE/C/common/inc/cutil_inline.h
#include <cutil_inline.h>

// =====================================================================================
// Device code
// =====================================================================================

__global__ void VecAdd(const float* gInputData, float* gOutputData, unsigned int nvec)
{
    int idx = nvec * blockDim.x * gridDim.x;
	idx += (blockDim.x * blockIdx.x + threadIdx.x);
	//gOutputData[idx] = 0.0f;
	float tmpv = gInputData[idx]+gOutputData[idx];
	//gOutputData[idx] = sin(tmpv);
	gOutputData[idx] = (tmpv);
	__syncthreads();
}

 __global__ void sinCompute(float* gOutputData, unsigned int maxloop)
{
	for(unsigned int loop=0; loop<maxloop; ++loop)
	{
		unsigned int idx = loop * blockDim.x * gridDim.x;
		idx += (blockDim.x * blockIdx.x + threadIdx.x);
		if(idx<100000000)
		{
			gOutputData[idx] = sin(sqrtf((float)(idx+1)));
			//gOutputData[idx] = __sinf(sqrtf((float)(idx+1)));
		}
		else
		{
			break;
		}
	}
}

__global__ void writeZeroOutput(int* gOutputData, unsigned int datasize, unsigned int intcountperthread)
{
	unsigned int idxbase = (blockDim.x * blockIdx.x + threadIdx.x)*intcountperthread;
	for(unsigned int loop=0; loop<intcountperthread; ++loop)
	{
		unsigned int idx = loop + idxbase;
		if(idx<datasize)
		{
			gOutputData[idx] = 0;
		}
		else
		{
			break;
		}
	}
	
}

__global__ void makeFullbuildDoc_bak(char* gInputData, int* gOutputData, unsigned int idatasize, unsigned int datasizeperthread, unsigned int odatasize, unsigned int intcountperthread)
{	
	unsigned int idxbase = (blockDim.x * blockIdx.x + threadIdx.x)*datasizeperthread;
	unsigned int outidxbase = (blockDim.x * blockIdx.x + threadIdx.x)*intcountperthread;
	unsigned int wpos = 0;

	char ch;
	int kvLineCount = 0;		// 已经读取的kvdump的信息行数
	int kvValueCount = 0;		// 当前正在处理的信息字段号（在每个信息行内的）
	int docCount = 0;			// 已经处理的文档数

	//int finddocend;		// 循环处于已经找到末尾的状态
	//int finddocbegin;		// 循环处于已经找到头的状态
	
	int bufferfrom = 0;
	int bufferto = 0;
	unsigned int loop=0;

	unsigned int cploop=0;

	char nowofferid[32];
	nowofferid[0]='\0';

	int strsame = 1;
	unsigned int idx;

	// 先找到一个信息行开始的地方
	while(true)
	{
		idx = loop + idxbase;

		if(idx >= idatasize)
		{
			break;
		}
		ch = gInputData[idx];
		if(8 == ch)
		{
			//gOutputData[idxbase+wpos] = '<';++ wpos;
			break;
		}
		++loop;
	}
	
	unsigned int till_last_char0 = 0;
	while(true)
	{	
		if(loop>datasizeperthread)
		{
			//gOutputData[idxbase+wpos] = ']';++ wpos;
			break;
		}

		idx = loop + idxbase;
		//finddocend = 0;

		if(idx >= idatasize)
		{
			//gOutputData[idxbase+wpos] = '|';++wpos;
			break;
		}

		ch = gInputData[idx];
		
		if('\0' == ch)
		{
			
			if(0 == kvValueCount)
			{
				// 信息id
				if(0 < (bufferto - bufferfrom))
				{
					cploop=0;
					strsame = 1;
					while(cploop < (bufferto - bufferfrom))
					{
						if(cploop > 31)
						{
							strsame=0;
							break;
						}
						if(gInputData[bufferfrom+cploop] != nowofferid[cploop])
						{
							strsame=0;
							break;
						}
						++ cploop;
					}

					if(0 == strsame)
					{
						// 如果当前读到的offerid 和 最近处理的offerid不一致，表示一个新的doc的开始
						if(0 < docCount)
						{
							//finddocend = 1;
							/*
							gOutputData[idxbase+wpos] = '<';
							++wpos;
							gOutputData[idxbase+wpos] = '/';
							++wpos;
							gOutputData[idxbase+wpos] = 'd';
							++wpos;
							gOutputData[idxbase+wpos] = 'o';
							++wpos;
							gOutputData[idxbase+wpos] = 'c';
							++wpos;
							gOutputData[idxbase+wpos] = '>';
							++wpos;
							gOutputData[idxbase+wpos] = '\1';
							++wpos;
							*/
							
							if(loop > datasizeperthread)
							{
								//gOutputData[idxbase+wpos] = '[';++wpos;
								break;
							}
						}
						/*
						gOutputData[idxbase+wpos] = '<';
						++wpos;
						gOutputData[idxbase+wpos] = 'd';
						++wpos;
						gOutputData[idxbase+wpos] = 'o';
						++wpos;
						gOutputData[idxbase+wpos] = 'c';
						++wpos;
						gOutputData[idxbase+wpos] = '>';
						++wpos;
						gOutputData[idxbase+wpos] = '\1';
						++wpos;
						gOutputData[idxbase+wpos] = '\n';
						++wpos;
						gOutputData[idxbase+wpos] = 'i';
						++wpos;
						gOutputData[idxbase+wpos] = 'd';
						++wpos;
						gOutputData[idxbase+wpos] = '=';
						++wpos;
						*/

						cploop=0;
						//*
						while(cploop <= (bufferto - bufferfrom))
						{
							nowofferid[cploop] = gInputData[bufferfrom+cploop];
							//gOutputData[idxbase+wpos] = nowofferid[cploop];
							//++wpos;
							++ cploop;
						}
						//*/
						
						if(wpos >= intcountperthread)
						{
							break;
						}
						gOutputData[outidxbase+wpos] = -1;
						++wpos;

						if(wpos >= intcountperthread-1)
						{
								break;
						}
						gOutputData[outidxbase+wpos] = bufferfrom;
						++wpos;
						gOutputData[outidxbase+wpos] = bufferto;
						++wpos;
					
						++ docCount;
					}
				}
			}
			else if(1 == kvValueCount)
			{
				// 键名
				if(0 < (bufferto - bufferfrom))
				{
					if(wpos >= intcountperthread-1)
					{
						break;
					}
					gOutputData[outidxbase+wpos] = bufferfrom;
					++wpos;
					gOutputData[outidxbase+wpos] = bufferto;
					++wpos;
				}
			}
			else if(2 == kvValueCount)
			{
				if(0 == till_last_char0)
				{
					gOutputData[outidxbase+wpos] = -2;
					++wpos;
					gOutputData[outidxbase+wpos] = -2;
					++wpos;
				}
				else
				{
					// 键值
					if(0 <= (bufferto - bufferfrom))
					{
						if(wpos >= intcountperthread-1)
						{
							break;
						}
						gOutputData[outidxbase+wpos] = bufferfrom;
						++wpos;
						gOutputData[outidxbase+wpos] = bufferto;
						++wpos;
					}
				}
			}
			else
			{
			}
		
			till_last_char0 = 0;
			++ kvValueCount;
			bufferfrom = 0;
		}
		else if(8 == (int)ch)
		{
			++ till_last_char0;
			//gOutputData[idxbase+wpos] = '8';++ wpos;
			// ascii 8 是各行数据间的分隔符号
			++ kvLineCount;
			kvValueCount = 0;
			bufferfrom = 0;
		}
		else
		{
			++ till_last_char0;
			if(0 == bufferfrom)
			{
				bufferfrom = idx;
			}
			bufferto = idx;
		}
		++ loop;
	}
}

__global__ void makeBigTopHeap(int* gOutputData, unsigned int listsize)
{
}

__device__ float getRatioGpu(const unsigned int slen, const unsigned int tlen, const unsigned int sv, const unsigned int tv)
{
	if(slen == tlen)
	{
		// y方向不变 s == t
		return 1.0f;
	}
	
	float inratio = 0.0f;
	const float stradtio = float(slen)/float(tlen);

	if(stradtio<1.0f)
	{
		// t > s, 源像素投影在目标图像上源像素框比目标像素框较大
		if((float)sv <= float(tv)*stradtio)
		{
			// 源像素上边缘 比 目标像素上边缘 小
			if((float)(sv+1) < float(tv)*stradtio)
			{
				// 源像素下边缘在目标像素上边缘以上
				return 0.0f;
			}
			else
			{
				if((float)(sv+1) < float(tv+1)*stradtio)
				{
					// 源像素下边缘在目标像素上下边缘之间
					inratio = float(sv+1)/stradtio - tv;
				}
				else
				{
					// 源像素下边缘在目标像素下边缘以下
					inratio = 1.0f;
				}
			}
		}
		else if((float)sv <= float(tv+1)*stradtio)
		{
			// 源像素上边缘在目标像素上下边缘之间
			if((float)(sv+1) < float(tv+1)*stradtio)
			{
				// 源像素下边缘在目标像素下边缘以上
				return 0.0f;
			}
			else
			{
				// 源像素下边缘在目标像素下边缘以下
				inratio = float(tv+1)-(float(sv)/stradtio);
			}
		}
		else
		{
			// 源像素上边缘在目标像素下边缘以下
			return 0.0f;
		}
	}
	else
	{
		// y方向缩小，t < s, 源像素投影在目标图像上源像素框比目标像素框较小
		if((float)sv <= float(tv)*stradtio)
		{
			// 源像素上边缘在目标像素上边缘以上
			if((float)(sv+1) < float(tv)*stradtio)
			{
				// 源像素下边缘在目标像素上边缘以上
				return 0.0f;
			}
			else
			{
				if((float)(sv+1) < float(tv+1)*stradtio)
				{
					// 源像素下边缘在目标像素下边缘以上
					inratio = float(sv+1)-((float)(tv)*stradtio);//sw * ((sx+1/sw) - (tx/tw)) 
				}
				else
				{
					// 源像素下边缘在目标像素下边缘以下
					return 0.0f;
				}
			}
		}
		else if((float)sv <= float(tv+1)*stradtio)
		{
			// 源像素上边缘在目标像素下边缘以上
			if((float)(sv+1) < float(tv+1)*stradtio)
			{
				// 源像素下边缘在目标像素下边缘以上
				inratio = 1.0f;
			}
			else
			{
				// 源像素下边缘在目标像素下边缘以下
				inratio = ((float)(tv+1)*stradtio) - (float)sv; //sw * (tx+1/tw - sx/sw)
			}
		}
		else
		{
			// 源像素上边缘在目标像素下边缘以下
			return 0.0f;
		}

		inratio /= stradtio;
	}
	return inratio;
}

/* 
	gInputData		源image数据
	gOutputData		目标image数据
	swidth sheight	源image宽和高
	twidth theight	目标image宽和高
	bytesperpix		每个像素的色深byte数
	lineAlignByte	每行像素为了对齐，多出来的byte数
 */
__global__ void zoomImageGpu(
	unsigned char* gInputData, unsigned char* gOutputData, 
	const unsigned int swidth, const unsigned int sheight,
	const unsigned int twidth, const unsigned int theight,
	const unsigned int bytesperpix, const unsigned int lineAlignByte)
{
	
	unsigned int ty=blockIdx.x;
	while(true)
	{
		if(ty>=theight)
		{
			break;;
		}
		unsigned int tx=threadIdx.x;
		while(true)
		{
			if(tx>=twidth)
			{
				break;
			}
			const float stwradtio = float(swidth)/float(twidth);	// 源image和目标image 的宽度比
			const float sthradtio = float(sheight)/float(theight);	// 源image和目标image 的高度比
			//计算四个顶点投影到源image的坐标位置（lu，ru，lb，rb）
			unsigned int sxfrom, sxto, syfrom, syto;
			sxfrom = (unsigned int)(float(tx)*stwradtio); 
			sxfrom = sxfrom>=swidth ? swidth-1 : sxfrom;
			sxto = (unsigned int)(float(tx+1)*stwradtio); 
			sxto = sxto>=swidth ? swidth-1 : sxto;
			syfrom = (unsigned int)(float(ty)*sthradtio); 
			syfrom = syfrom>=sheight ? sheight-1 : syfrom;
			syto = (unsigned int)(float(ty+1)*sthradtio); 
			syto = syto>=sheight ? sheight-1 : syto;

			float sumv[4];
			float ratiosum = 0.0f;
			float sxinratio = 0.0f;
			float syinratio = 0.0f;
			float rgbratio = 0.0f;

			for(unsigned int i=0;i<bytesperpix;++i)
			{
				sumv[i] = 0.0f;
			}

			unsigned int sy=0;
			unsigned int sx=0;
			unsigned char* srgbpos;
			for(sy=syfrom; sy<=syto; ++sy)
			{
				syinratio = getRatioGpu(sheight, theight, sy, ty);
				for(sx=sxfrom; sx<=sxto; ++sx)
				{
					//foreach 源image,X坐标在[lu.x, ru.x]并且Y坐标在[lu.y, rb.y]的每个像素simg
					// 计算该在源image上的像素在目标image上投影的面积，占目标像素的面积比值 K(0.0f, 1.0f]
					// 源像素投影在目标像素上
					sxinratio  = getRatioGpu(swidth, twidth, sx, tx);
					rgbratio = sxinratio*syinratio;
					ratiosum += rgbratio;

					srgbpos = gInputData;
					srgbpos += ((sy*swidth+sx)*bytesperpix)+sy*lineAlignByte;
					for(unsigned int i=0;i<bytesperpix;++i)
					{
						sumv[i] += rgbratio*(float)srgbpos[i];
					}
				}
			}			
			unsigned char* trgbpos = gOutputData;
			trgbpos += (ty*twidth+tx)*bytesperpix;
			for(unsigned int i=0;i<bytesperpix;++i)
			{
				trgbpos[i] = (char)roundf(sumv[i]);
			}
			tx+=512;
		}
		ty+=512;
	}
}

// =====================================================================================
//Host code
// =====================================================================================

// 检查GPU参数
int gpuStatQuery(int argc, char** argv)
{
	// 获取显示设备个数
	int devcount;
	CUDA_SAFE_CALL(cudaGetDeviceCount(&devcount));
	printf("in gpuStatQuery, devcount=%d\n", devcount);

	if(devcount<1)
	{
		printf("in gpuStatQuery, no cuda gpu found\n");
		return 0;
	}

	// 获取显示设备参数
	// cudaDeviceProp define in /usr/local/cuda/include/driver_types.h
	/*
	char   name[256];                 ///< ASCII string identifying device
	size_t totalGlobalMem;            ///< Global memory available on device in bytes
	size_t sharedMemPerBlock;         ///< Shared memory available per block in bytes
	int    regsPerBlock;              ///< 32-bit registers available per block
	int    warpSize;                  ///< Warp size in threads
	size_t memPitch;                  ///< Maximum pitch in bytes allowed by memory copies
	int    maxThreadsPerBlock;        ///< Maximum number of threads per block
	int    maxThreadsDim[3];          ///< Maximum size of each dimension of a block
	int    maxGridSize[3];            ///< Maximum size of each dimension of a grid
	int    clockRate;                 ///< Clock frequency in kilohertz
	size_t totalConstMem;             ///< Constant memory available on device in bytes
	int    major;                     ///< Major compute capability
	int    minor;                     ///< Minor compute capability
	size_t textureAlignment;          ///< Alignment requirement for textures
	int    deviceOverlap;             ///< Device can concurrently copy memory and execute a kernel
	int    multiProcessorCount;       ///< Number of multiprocessors on device
	int    kernelExecTimeoutEnabled;  ///< Specified whether there is a run time limit on kernels
	int    integrated;                ///< Device is integrated as opposed to discrete
	int    canMapHostMemory;          ///< Device can map host memory with cudaHostAlloc/cudaHostGetDevicePointer
	int    computeMode;               ///< Compute mode (See ::cudaComputeMode)
	int    maxTexture1D;              ///< Maximum 1D texture size
	int    maxTexture2D[2];           ///< Maximum 2D texture dimensions
	int    maxTexture3D[3];           ///< Maximum 3D texture dimensions
	int    maxTexture2DArray[3];      ///< Maximum 2D texture array dimensions
	size_t surfaceAlignment;          ///< Alignment requirements for surfaces
	int    concurrentKernels;         ///< Device can possibly execute multiple kernels concurrently
	int    ECCEnabled;                ///< Device has ECC support enabled
	int    pciBusID;                  ///< PCI bus ID of the device
	int    pciDeviceID;               ///< PCI device ID of the device
	int    __cudaReserved[22];
	*/

	for(int i=0; i< devcount; ++i)
	{
		printf("\n=============== check gpu %d ===============\n",i);
		cudaDeviceProp deviceProp;

		CUDA_SAFE_CALL(cudaGetDeviceProperties(&deviceProp, i));
		CUT_CHECK_ERROR("fail to call cudaGetDeviceProperties");

		printf("deviceProp.name=%s\n",				deviceProp.name);
		printf("deviceProp.totalGlobalMem=%d\n",	deviceProp.totalGlobalMem);
		printf("deviceProp.sharedMemPerBlock=%d\n", deviceProp.sharedMemPerBlock);
		printf("deviceProp.regsPerBlock=%d\n",		deviceProp.regsPerBlock);
		printf("deviceProp.warpSize=%d\n",			deviceProp.warpSize);
		printf("deviceProp.memPitch=%d\n",			deviceProp.memPitch);
		printf("deviceProp.maxThreadsPerBlock=%d\n", deviceProp.maxThreadsPerBlock);
		printf("deviceProp.maxThreadsDim=[%d,%d,%d]\n",		deviceProp.maxThreadsDim[0], deviceProp.maxThreadsDim[1], deviceProp.maxThreadsDim[2]);
		printf("deviceProp.maxGridSize=[%d,%d,%d]\n",		deviceProp.maxGridSize[0],deviceProp.maxGridSize[1], deviceProp.maxGridSize[2]);
		printf("deviceProp.clockRate=%d\n",			deviceProp.clockRate);
		printf("deviceProp.totalConstMem=%d\n",		deviceProp.totalConstMem);
		printf("deviceProp.major=%d\n",				deviceProp.major);
		printf("deviceProp.minor=%d\n",				deviceProp.minor);
		printf("deviceProp.textureAlignment=%d\n",	deviceProp.textureAlignment);
		printf("deviceProp.deviceOverlap=%d\n",		deviceProp.deviceOverlap);
		printf("deviceProp.multiProcessorCount=%d\n",		deviceProp.multiProcessorCount);
		printf("deviceProp.kernelExecTimeoutEnabled=%d\n",	deviceProp.kernelExecTimeoutEnabled);
		printf("deviceProp.integrated=%d\n",		deviceProp.integrated);
		printf("deviceProp.canMapHostMemory=%d\n",	deviceProp.canMapHostMemory);
		printf("deviceProp.computeMode=%d\n",		deviceProp.computeMode);
		printf("deviceProp.maxTexture1D=%d\n",		deviceProp.maxTexture1D);
		printf("deviceProp.maxTexture2D=[%d,%d]\n", deviceProp.maxTexture2D[0], deviceProp.maxTexture2D[1]);
		printf("deviceProp.maxTexture3D=[%d.%d.%d]\n",		deviceProp.maxTexture3D[0],deviceProp.maxTexture3D[1],deviceProp.maxTexture3D[2]);
		printf("deviceProp.maxTexture2DArray[%d,%d,%d]\n",	deviceProp.maxTexture2DArray[0],deviceProp.maxTexture2DArray[1],deviceProp.maxTexture2DArray[2]);
		printf("deviceProp.surfaceAlignment=%d\n",	deviceProp.surfaceAlignment);
		printf("deviceProp.concurrentKernels=%d\n", deviceProp.concurrentKernels);
		printf("deviceProp.ECCEnabled=%d\n",		deviceProp.ECCEnabled);
		printf("deviceProp.pciBusID=%d\n",			deviceProp.pciBusID);
		printf("deviceProp.pciDeviceID=%d\n",		deviceProp.pciDeviceID);
	}
	return 1;
}

// vector add test
int vectorAddTest(int argc, char** argv)
{
	printf("\n================ in vectorAddTest ================\n");
	const unsigned int nblock = 8;
	const unsigned int nthread = 8;
	const unsigned nvec = 100000;

	const unsigned int memsize = sizeof(float)*nblock*nthread*nvec;
	float * hostInputData = (float*)malloc(memsize);
	float * hostOutputData = (float*)malloc(memsize);

	for(unsigned int tmpv=0; tmpv<nvec; ++tmpv)
	{
		for(unsigned int loop=0;loop<nblock*nthread;++loop)
		{
			hostInputData[tmpv*nblock*nthread + loop] = loop;
			hostOutputData[tmpv*nblock*nthread + loop] = 0.0f;
		}
	}
	for(unsigned int loop=0;loop<nblock*nthread;++loop)
	{
		printf("in [%d] = %f\n", loop, hostInputData[loop]);
	}

	clock_t nowt1, nowt2, nowt3, nowt4, nowt5;
	nowt1 = clock();
	// 初始化CUDA环境
	CUT_DEVICE_INIT(argc, argv);
	CUT_CHECK_ERROR("fail to call CUT_DEVICE_INIT");

	// 分配设备内存存放输入数据
	float* deviceInputData;
	CUDA_SAFE_CALL(cudaMalloc((void**)&deviceInputData, memsize));
	CUT_CHECK_ERROR("fail to call cudaMalloc for deviceInputData");

	// 分配设备内存存放输出数据
	float* deviceOutputData;
	CUDA_SAFE_CALL(cudaMalloc((void**)&deviceOutputData, memsize));
	CUT_CHECK_ERROR("fail to call cudaMalloc for deviceOutputData");

	// 拷贝主机数据到设备
	CUDA_SAFE_CALL(cudaMemcpy(deviceInputData, hostInputData, memsize, cudaMemcpyHostToDevice));
	CUT_CHECK_ERROR("fail to call cudaMemcpy to copy hostInputData to deviceInputData");
	CUDA_SAFE_CALL(cudaMemcpy(deviceOutputData, hostOutputData, memsize, cudaMemcpyHostToDevice));
	CUT_CHECK_ERROR("fail to call cudaMemcpy to copy hostOutputData to deviceOutputData");
	
	dim3 testgrid(nblock, 1, 1);
	dim3 testthread(nthread, 1, 1);
	
	nowt2 = clock();
	// 运行内核函数
	for(unsigned int tmpv=0; tmpv<nvec; ++tmpv)
	{
		VecAdd<<<testgrid, testthread>>>(deviceInputData, deviceOutputData, tmpv);
		CUT_CHECK_ERROR("fail to call kernal function");
	}
	nowt3 = clock();
	nowt4 = clock();

	// 拷贝设备数据到主机
	CUDA_SAFE_CALL(cudaMemcpy(hostOutputData, deviceOutputData, memsize, cudaMemcpyDeviceToHost));
	CUT_CHECK_ERROR("fail to call cudaMemcpy to copy deviceOutputData to hostOutputData");

	// 释放申请的设备端内存
	CUDA_SAFE_CALL(cudaFree(deviceInputData));
	CUDA_SAFE_CALL(cudaFree(deviceOutputData));

	nowt5 = clock();
	for(unsigned int loop=0;loop<nblock*nthread;++loop)
	{
		printf("out [%d] = %f\n", loop, hostOutputData[loop]);
	}
	printf("\nclockuse: copy_to_device=%d, compute1=%d, compute2=%d, fetch_from_device=%d\n",(int)(nowt2-nowt1),(int)(nowt3-nowt2),(int)(nowt4-nowt3),(int)(nowt5-nowt4));
	printf("CLOCKS_PER_SEC=%ld \n",CLOCKS_PER_SEC);
	// 退出CUDA环境
	CUT_EXIT(argc, argv);
	CUT_CHECK_ERROR("fail to call CUT_EXIT");

	return 0;
}


/*
100,000,000（1亿）个浮点数求取平方根后的sin值(sin(sqrt(x))), as fast as you can。

1） 简单起见，就从1开始到100,000,000依次填充到输入。
2） 请使用你所知道的所有手段（什么多进程，多线程，内存映射，甚至网络集群，能用上的都用上）在最短的时间完成计算，将计算结果填充到内存。
3） 请记录从程序开始到100,000,000个结果填充到内存所消耗的时间，精确到ms。
4） 最后将结果写入磁盘，每行一个，保留3位小数，用以检查程序的正确性，结果写入磁盘的时间不计算。
*/
int sinComputeTest(int argc, char** argv)
{
	timeval clkBegin;
	gettimeofday(&clkBegin, NULL);

	//printf("\n================ in vectorAddTest ================\n");
	const unsigned int nblock = 100;
	const unsigned int nthread = 100;
	const unsigned nvec = 10000;

	const unsigned int memsize = sizeof(float)*nblock*nthread*nvec;
	float * hostOutputData = (float*)malloc(memsize);

	// 初始化CUDA环境
	CUT_DEVICE_INIT(argc, argv);
	CUT_CHECK_ERROR("fail to call CUT_DEVICE_INIT");

	// 分配设备内存存放输出数据
	float* deviceOutputData;
	CUDA_SAFE_CALL(cudaMalloc((void**)&deviceOutputData, memsize));
	CUT_CHECK_ERROR("fail to call cudaMalloc for deviceOutputData");

	timeval clkMalloc;
	gettimeofday(&clkMalloc, NULL);

	// 运行内核函数
	dim3 testgrid(128, 1, 1);
	dim3 testthread(128, 1, 1);
	unsigned int maxloop = (unsigned int)((100000000/(double)testgrid.x)/testthread.x);
	maxloop += 1;
	printf("maxloop: %d\n", maxloop);
	//for(unsigned int tmpv=0; tmpv<nvec; ++tmpv)
	{
		sinCompute<<<testgrid, testthread>>>(deviceOutputData, maxloop);
		CUT_CHECK_ERROR("fail to call kernal function");
	}
	
	timeval clkGpu;
	gettimeofday(&clkGpu, NULL);

	// 拷贝设备数据到主机
	CUDA_SAFE_CALL(cudaMemcpy(hostOutputData, deviceOutputData, memsize, cudaMemcpyDeviceToHost));
	CUT_CHECK_ERROR("fail to call cudaMemcpy to copy deviceOutputData to hostOutputData");

	timeval clkEnd;
	gettimeofday(&clkEnd, NULL);

	// 释放申请的设备端内存
	CUDA_SAFE_CALL(cudaFree(deviceOutputData));

	// 打印时间信息
	long timemalcusec = (clkMalloc.tv_usec - clkBegin.tv_usec);
	timemalcusec += 1000000L*(clkMalloc.tv_sec - clkBegin.tv_sec);
	long timecostusec = (clkEnd.tv_usec - clkBegin.tv_usec);
	timecostusec += 1000000L*(clkEnd.tv_sec - clkBegin.tv_sec);
	
	long timegpuusec = clkGpu.tv_usec - clkMalloc.tv_usec;
	timegpuusec += 1000000L*(clkGpu.tv_sec - clkMalloc.tv_sec);
	long timecpresusec = clkEnd.tv_usec - clkGpu.tv_usec;
	timecpresusec += 1000000L*(clkEnd.tv_sec - clkGpu.tv_sec);

	long timecostMs = timecostusec/1000L;

	printf("malloc time: %ld us\n", timemalcusec);
	printf("compute time: %ld us\n", timecostusec);
	printf("gpu time: %ld us\n", timegpuusec);
	printf("copy res to mem time: %ld us\n", timecpresusec);
	printf("compute time: %ld ms\n", timecostMs);

	// 计算结果写入磁盘
	printf("write result to file: cpre.data");
	FILE* cpref = fopen("cpre.data", "w");
	int numv = nblock*nthread*nvec;
	//int numv = nblock*nthread;
	for(int cpt = 1; cpt<= numv; ++cpt)
	{
		fprintf(cpref, "%d: %.6f\n", cpt, hostOutputData[cpt-1]);
		if(cpt > 1000)
		{	
			break;
		}
	}
	fclose(cpref);

	timeval clkWrEnd;
	gettimeofday(&clkWrEnd, NULL);
	long timewrusec = clkWrEnd.tv_usec - clkEnd.tv_usec;
	timewrusec += 1000000L*(clkWrEnd.tv_sec - clkEnd.tv_sec);
	printf("write result to file: cpre.data, done, %ld us, %ld ms", timewrusec, (timewrusec/1000));


	free(hostOutputData);
	// 退出CUDA环境
	CUT_EXIT(argc, argv);
	CUT_CHECK_ERROR("fail to call CUT_EXIT");

	return 0;
}

/*
	从kvdb中dump出来的文件格式描述如下:
		由数据库表每行记录信息组成
		每行信息对一个字符串，不分行，信息间的分割符号是ascii码 8;
		每行记录信息对应的字符串组成如下：信息id, 信息键名, 信息键值, 键值插入时间文本, 键值插入时间戳, 补齐位
		行信息内部各字段值间的分割字符串为ascii码 0

		信息id相同的多行记录描述了一个完整的对象	


	fullbuild doc文件格式
		对象以<doc>开始，</doc>结束，开始和结束标示符号都单独占1行(\n)
		对象的每一个键值对占用1行，键名=键值
		每一行的末尾都有ascii码 1 结束(即 \01\n )
		对象的第一个键值对是id=idvalue

	计算从程序开始，到fullbuilddoc写入结果文件的时间
*/

//const char* srcFilePath = "/media/DATA/forCuda/kvdbdata/10.252.42.117-t1019_2";
const char* srcFilePath = "/home/mengchang/download/10.252.40.205-t1038_2";
//const char* disFilePath = "/home/jiening/testres.xml";
const char* disFilePath = "/home/mengchang/download/testres.xml";

int simuKVdumpToFullbuildDoc(int argc, char** argv)
{
	timeval clkBegin;
	gettimeofday(&clkBegin, NULL);
	
	FILE* fp = fopen(srcFilePath, "r");
	if(NULL == fp)
	{
		printf("in simuKVdumpToFullbuildDoc, fail to read source kvdb dump file, program over\n");
		return 0;
	}
	fseek(fp, 0L, SEEK_END);
	long fileByteCount = ftell(fp);
	char* freadbuffer = (char*)malloc(fileByteCount);
	fseek(fp, 0L, SEEK_SET);
	size_t readbytecount = fread(freadbuffer, 1, (size_t)fileByteCount, fp);
	if((size_t)fileByteCount != readbytecount)
	{
		printf("in simuKVdumpToFullbuildDoc, fread fail to read %ld bytes but %d byte\n", fileByteCount, readbytecount);
		return 0;
	}
	fclose(fp);
	timeval clkFileOpen;
	gettimeofday(&clkFileOpen, NULL);

	dim3 testgrid(32, 1, 1);
	dim3 testthread(128, 1, 1);

	int intcountperthread = fileByteCount/10000;
	intcountperthread /= testgrid.x;
	intcountperthread /= testthread.x;
	++intcountperthread;
	intcountperthread *= 200;

	int* deviceOutputData;
	unsigned int outDataSize = (unsigned int)(testgrid.x*testthread.x*intcountperthread);

	printf("intcountperthread: %d %x , fwritebuffer size %d\n", intcountperthread, intcountperthread, outDataSize*sizeof(int));

	int* fwritebuffer = (int*)malloc(outDataSize*sizeof(int));
	// 初始化CUDA环境
	CUT_DEVICE_INIT(argc, argv);
	CUT_CHECK_ERROR("fail to call CUT_DEVICE_INIT");

	
	// 分配设备内存存放输出数据
	char* deviceInputData;
	CUDA_SAFE_CALL(cudaMalloc((void**)&deviceInputData, fileByteCount));
	CUT_CHECK_ERROR("fail to call cudaMalloc for deviceInputData");

	CUDA_SAFE_CALL(cudaMalloc((void**)&deviceOutputData, outDataSize*sizeof(int)));
	CUT_CHECK_ERROR("fail to call cudaMalloc for deviceOutputData");

	// 拷贝主机数据到设备
	CUDA_SAFE_CALL(cudaMemcpy(deviceInputData, freadbuffer, fileByteCount, cudaMemcpyHostToDevice));
	CUT_CHECK_ERROR("fail to call cudaMemcpy to copy hostInputData to deviceInputData");
	timeval clkMalloc;
	gettimeofday(&clkMalloc, NULL);

	unsigned int datasizeperthread = (unsigned int)((fileByteCount/(double)testgrid.x)/testthread.x);
	datasizeperthread+= 1;
	printf("datasizeperthread: %d %x \n", datasizeperthread, datasizeperthread);
	{
		writeZeroOutput<<<testgrid, testthread>>>(deviceOutputData, (unsigned int)(testgrid.x*testthread.x*intcountperthread), intcountperthread);
		CUT_CHECK_ERROR("fail to call kernal function");
	}
	{
		makeFullbuildDoc_bak<<<testgrid, testthread>>>(deviceInputData, deviceOutputData, (unsigned int)fileByteCount, datasizeperthread, (unsigned int)(testgrid.x*testthread.x*intcountperthread), intcountperthread);
		CUT_CHECK_ERROR("fail to call kernal function");
	}
	
	timeval clkGpu;
	gettimeofday(&clkGpu, NULL);

	long timecostgpu = (clkGpu.tv_usec - clkMalloc.tv_usec);
	timecostgpu += 1000000L*(clkGpu.tv_sec - clkMalloc.tv_sec);
	printf("\ngpucompute:\t\t %ld us, %ld ms\n", timecostgpu, (timecostgpu/1000L));

	// 拷贝设备数据到主机
	CUDA_SAFE_CALL(cudaMemcpy(fwritebuffer, deviceOutputData, (outDataSize*sizeof(int)), cudaMemcpyDeviceToHost));
	CUT_CHECK_ERROR("fail to call cudaMemcpy to copy deviceOutputData to hostOutputData");

	// 释放申请的设备端内存
	CUDA_SAFE_CALL(cudaFree(deviceInputData));
	CUDA_SAFE_CALL(cudaFree(deviceOutputData));
	
	timeval clkCopytohost;
	gettimeofday(&clkCopytohost, NULL);

	// 计算结果写入文件
	FILE* wfp = fopen(disFilePath, "w");
	if(NULL == wfp)
	{
		printf("---bug: fail to open write file\n");
	}
	else
	{
		int intcount = 0;
		int valueReadCount = 0;
		int doccount = 0;
		for(int i=0; i < outDataSize; ++i)
		{
			if(0 == fwritebuffer[i])
			{
				intcount = 0;
			}
			else
			{
				if(-1 == fwritebuffer[i])
				{
					fprintf(wfp, "</doc>\n<doc>\nid=");
					intcount = 0;
					valueReadCount = 1;
					++ doccount;
				}
				else
				{
					++intcount;
					//fprintf(wfp, "\n%d, %d\n", i, fwritebuffer[i]);
					if(0 == intcount%2)
					{
						++ valueReadCount;
						
						int valuesize = fwritebuffer[i]-fwritebuffer[i-1];
						if(-2 == fwritebuffer[i] && -2 == fwritebuffer[i-1])
						{
							fprintf(wfp, "---find empty value\n");
						}
						else
						{
							//fprintf(wfp, "%d, value size %d\n", i, valuesize);
							if(0 <= valuesize)
							{
								if(valuesize > 100000)
								{
									fprintf(wfp, "---bug here\n");
								}
								else
								{
									char* buferpos = freadbuffer;
									buferpos += fwritebuffer[i-1];

									fwrite(buferpos, 1, (valuesize+1), wfp);
									if(1 == valueReadCount%2)
									{
										fwrite("=", 1, (1), wfp);
									}
									else
									{
										fwrite("\n", 1, (2), wfp);
									}
								}
							}
							else
							{
								fprintf(wfp, "---bug here\n");
							}
						}
					}
				}
			}
			//if(i>100000)break;
		}
		printf("doccount=%d\n", doccount);

		fclose(wfp);
	}
	timeval clkEnd;
	gettimeofday(&clkEnd, NULL);

	// 打印时间 
	printf("\nfileByteCount = %ld\n", fileByteCount);

	long timecostfopen = (clkFileOpen.tv_usec - clkBegin.tv_usec);	
	timecostfopen += 1000000L*(clkFileOpen.tv_sec - clkBegin.tv_sec);	
	printf("\nfileload:\t\t %ld us, %ld ms\n", timecostfopen, (timecostfopen/1000L));

	long timecostmalloc = (clkMalloc.tv_usec -  clkFileOpen.tv_usec);
	timecostmalloc += 1000000L*(clkMalloc.tv_sec -  clkFileOpen.tv_sec);
	printf("\nmallocgpu:\t\t %ld us, %ld ms\n", timecostmalloc, (timecostmalloc/1000L));

	//long timecostgpu = (clkGpu.tv_usec - clkMalloc.tv_usec);
	//timecostgpu += 1000000L*(clkGpu.tv_sec - clkMalloc.tv_sec);
	printf("\ngpucompute:\t\t %ld us, %ld ms\n", timecostgpu, (timecostgpu/1000L));

	long timecostcopyhost = (clkCopytohost.tv_usec - clkGpu.tv_usec);
	timecostcopyhost += 1000000L*(clkCopytohost.tv_sec - clkGpu.tv_sec);
	printf("\ngpu2mem:\t\t %ld us, %ld ms\n", timecostcopyhost, (timecostcopyhost/1000L));

	long timecostfwrite = (clkEnd.tv_usec - clkCopytohost.tv_usec);
	timecostfwrite += 1000000L*(clkEnd.tv_sec - clkCopytohost.tv_sec);
	printf("\nfilewrite:\t\t %ld us, %ld ms\n", timecostfwrite, (timecostfwrite/1000L));

	long timecostusec = (clkEnd.tv_usec - clkBegin.tv_usec);
	timecostusec += 1000000L*(clkEnd.tv_sec - clkBegin.tv_sec);
	printf("\ntotaltime:\t\t %ld us, %ld ms\n", timecostusec, (timecostusec/1000L));

	return 0;
}

int testImageZoom(int argc, char** argv)
{
	printf("\nin func testImageZoom\n");
	const unsigned int inputmemsize = 4*1024*4*1024*4;
	const unsigned int outputmemsize = 2*1024*2*1024*4;

	const unsigned int byteperpix = 3;
	//const unsigned int twidth = 320;
	//const unsigned int theight= 320;
	const unsigned int twidth = 1024;
	const unsigned int theight= 1024;
	const unsigned int outimgsize=twidth*theight*byteperpix;
	
	
	unsigned char* hostInputData;
	unsigned char* hostOutputData = (unsigned char*)malloc(outputmemsize);

	// ================= image read from file using freeimage start here =================
	/*
	// zoom in test1
	const char* sImgPath = "/home/mengchang/IMGP5870.JPG";
	const char* tImgPath = "/home/mengchang/test1.bmp";
	// zoom in test2
	const char* sImgPath = "/home/mengchang/Avatar_002009.jpg";
	const char* tImgPath = "/home/mengchang/test2.bmp";
	// zoom in test3
	const char* sImgPath = "/home/mengchang/acmod.jpg";
	const char* tImgPath = "/home/mengchang/test3.bmp";
	*/
	// zoom out test 1
	const char* sImgPath = "/home/mengchang/test3.JPG";
	const char* tImgPath = "/home/mengchang/test3big.bmp";


	FIBITMAP * dib = FreeImage_Load(FIF_JPEG, sImgPath, BMP_DEFAULT);

	if(NULL == dib)
	{
		printf("fail to read image Avatar_002009.jpg\n");
		return 0;
	}
	printf("succ to read image %s\n", sImgPath);
	const unsigned int bitsPerPix = FreeImage_GetBPP(dib);
	const unsigned int inImageWidth = FreeImage_GetWidth(dib);
	const unsigned int inImageHeight = FreeImage_GetHeight(dib);
	const unsigned int inImagePitch = FreeImage_GetPitch(dib);
	FREE_IMAGE_TYPE inImageType = FreeImage_GetImageType(dib);
	FREE_IMAGE_COLOR_TYPE inImageColorType = FreeImage_GetColorType(dib);

	const unsigned int linePixBytes = inImageWidth*bitsPerPix/8;
	const unsigned int inImageLineAlignByte = (inImagePitch > linePixBytes) ? inImagePitch - linePixBytes : 0;
	printf("inImageType=%d\nbitsPerPix=%d\ninImageWidth=%d\ninImageHeight=%d\ninImagePitch=%d\ninImageColorType=%d\ninImageLineAlignByte=%d\n\n", 
		inImageType,
		bitsPerPix,
		inImageWidth,
		inImageHeight,
		inImagePitch,
		inImageColorType,
		inImageLineAlignByte);
	// ================= image read from file using freeimage end here =================

	hostInputData = (unsigned char*)FreeImage_GetBits(dib);
	unsigned int swidth = inImageWidth;
	unsigned int sheight = inImageHeight;
	unsigned int fileByteCount = swidth*sheight*byteperpix;

	timeval clkBegin;
	gettimeofday(&clkBegin, NULL);

	// 初始化CUDA环境
	CUT_DEVICE_INIT(argc, argv);
	CUT_CHECK_ERROR("fail to call CUT_DEVICE_INIT");

	// 分配设备内存存放输出数据
	unsigned char* deviceOutputData;
	CUDA_SAFE_CALL(cudaMalloc((void**)&deviceOutputData, outputmemsize));
	CUT_CHECK_ERROR("fail to call cudaMalloc for deviceOutputData");

	unsigned char* deviceInputData;
	CUDA_SAFE_CALL(cudaMalloc((void**)&deviceInputData, inputmemsize));
	CUT_CHECK_ERROR("fail to call cudaMalloc for deviceInputData");
	timeval clkMalloc;
	gettimeofday(&clkMalloc, NULL);

	CUDA_SAFE_CALL(cudaMemcpy(deviceInputData, hostInputData, fileByteCount, cudaMemcpyHostToDevice));
	CUT_CHECK_ERROR("fail to call cudaMemcpy to copy hostInputData to deviceInputData");
	timeval clkCpy2device;
	gettimeofday(&clkCpy2device, NULL);

	// 运行内核函数
	dim3 testgrid(theight>512?512:theight, 1, 1);
	dim3 testthread(twidth>512?512:theight, 1, 1);
	{
		zoomImageGpu<<<testgrid, testthread>>>(
			deviceInputData, deviceOutputData, 
			swidth, sheight, twidth, theight, 
			byteperpix, inImageLineAlignByte);

		CUT_CHECK_ERROR("fail to call kernal function");
	}
	
	timeval clkGpu;
	gettimeofday(&clkGpu, NULL);

	// 拷贝设备数据到主机
	CUDA_SAFE_CALL(cudaMemcpy(hostOutputData, deviceOutputData, outimgsize, cudaMemcpyDeviceToHost));
	CUT_CHECK_ERROR("fail to call cudaMemcpy to copy deviceOutputData to hostOutputData");

	timeval clkEnd;
	gettimeofday(&clkEnd, NULL);

	// 释放申请的设备端内存
	CUDA_SAFE_CALL(cudaFree(deviceInputData));
	CUDA_SAFE_CALL(cudaFree(deviceOutputData));

	FIBITMAP * outImageBmp = FreeImage_AllocateT(FIT_BITMAP, twidth, theight, byteperpix*8);
	BYTE *outbits = (BYTE*)FreeImage_GetBits(outImageBmp);
	memcpy((char*)outbits, (char*)hostOutputData, outimgsize);

	if (FreeImage_Save(FIF_BMP, outImageBmp, tImgPath, BMP_DEFAULT))
	{
		// bitmap successfully saved!
		printf("succ to write image file %s\n", tImgPath);
	}
	else
	{
		printf("fail to write image file %s\n", tImgPath);
	}


	// 打印时间信息
	long timemalcusec = (clkMalloc.tv_usec - clkBegin.tv_usec);
	timemalcusec += 1000000L*(clkMalloc.tv_sec - clkBegin.tv_sec);

	long timecpy2devusec = (clkCpy2device.tv_usec - clkMalloc.tv_usec);
	timecpy2devusec += 1000000L*(clkCpy2device.tv_sec - clkMalloc.tv_sec);

	long timegpuusec = clkGpu.tv_usec - clkCpy2device.tv_usec;
	timegpuusec += 1000000L*(clkGpu.tv_sec - clkCpy2device.tv_sec);

	long timecpy2hostusec = clkEnd.tv_usec - clkGpu.tv_usec;
	timecpy2hostusec += 1000000L*(clkEnd.tv_sec - clkGpu.tv_sec);

	long timecostusec = (clkEnd.tv_usec - clkBegin.tv_usec);
	timecostusec += 1000000L*(clkEnd.tv_sec - clkBegin.tv_sec);

	long timecostMs = timecostusec/1000L;

	printf("malloc time: %ld us\n", timemalcusec);
	printf("cpy2dev time: %ld us\n", timecpy2devusec);
	printf("gpu time: %ld us\n", timegpuusec);
	printf("cpy2host time: %ld us\n", timecpy2hostusec);
	printf("gpu deal time: %ld us\n", timecpy2devusec+timegpuusec+timecpy2hostusec);
	
	printf("compute time: %ld us\n", timecostusec);
	printf("compute time: %ld ms\n", timecostMs);

	return 0;
}

int main(int argc, char** argv)
{
	printf("in jieningTest.cu size of float=%d \n", sizeof(float));

	// 检查GPU参数
	if(0 == gpuStatQuery(argc, argv))
	{
		return 0;
	}

	// 向量相加
	//vectorAddTest(argc, argv);

	// sin值计算
	//sinComputeTest(argc, argv);

	// kvdump to fullbuild doc
	//printf("\n\nin jieningTest.cu to call simuKVdumpToFullbuildDoc\n\n");
	//simuKVdumpToFullbuildDoc(argc, argv);

	// image zoom in
	testImageZoom(argc, argv);
	return 0;
}
